/*
 * Software License Agreement (BSD License)
 *
 *  Point Cloud Library (PCL) - www.pointclouds.org
 *  Copyright (c) 2011, Willow Garage, Inc.
 * 
 *  All rights reserved.
 *
 *  Redistribution and use in source and binary forms, with or without
 *  modification, are permitted provided that the following conditions
 *  are met:
 *
 *   * Redistributions of source code must retain the above copyright
 *     notice, this list of conditions and the following disclaimer.
 *   * Redistributions in binary form must reproduce the above
 *     copyright notice, this list of conditions and the following
 *     disclaimer in the documentation and/or other materials provided
 *     with the distribution.
 *   * Neither the name of Willow Garage, Inc. nor the names of its
 *     contributors may be used to endorse or promote products derived
 *     from this software without specific prior written permission.
 *
 *  THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
 *  "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
 *  LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
 *  FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
 *  COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
 *  INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
 *  BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 *  LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
 *  CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
 *  LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
 *  ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
 *  POSSIBILITY OF SUCH DAMAGE.
 *
 */

#include "device.hpp"
#include <pcl/gpu/utils/device/funcattrib.hpp>
#include <pcl/gpu/utils/device/block.hpp>
#include <pcl/gpu/utils/device/warp.hpp>

#include "cuPrintf.cu"
#include <pcl/common/time.h>
#include "pointer_shift.cu"


typedef pcl::ScopeTime ScopeTimeT;


__global__ void
loadTsdfValueKernel(short2* startPtr, float4 *data_ptr, int size, const int divisor)
{
  
  int i = blockDim.x * blockIdx.x  + threadIdx.x;
  if(i < size)
  {  
    float4 *point_ptr = data_ptr + (i * 2);
    float4 *tsdfvalue_ptr = point_ptr + 1;  
  
    if( //Indices x,y,z are within [0...VOLUME_{X,Y,Z}]
        ( ((int)((*point_ptr).x) < pcl::device::VOLUME_X) && ((int)((*point_ptr).x) >= 0) ) || 
        ( ((int)((*point_ptr).y) < pcl::device::VOLUME_Y) && ((int)((*point_ptr).y) >= 0) ) ||
        ( ((int)((*point_ptr).z) < pcl::device::VOLUME_Z) && ((int)((*point_ptr).z) >= 0) )
      )
    {  
       short2 *dest_tsdf_index = startPtr + ( pcl::device::VOLUME_Y * pcl::device::VOLUME_X * (int)((*point_ptr).z) + ( (int) ( (*point_ptr).y) * pcl::device::VOLUME_X ) ) + (int)((*point_ptr).x);
       (*dest_tsdf_index).x = (*tsdfvalue_ptr).x * divisor;
       (*dest_tsdf_index).y = 1.0;
    }
  }
  else
    return; 
  
}

namespace pcl
{
  namespace device
  {

    void 
    pushCloudAsSliceGPU(const PtrStep<short2>& volume, pcl::gpu::DeviceArray<float4> cloud_gpu)
    {
   

      // Because every point is 8 floats.
      // X Y Z 0     
      // I 0 0 0
      int numPoints = cloud_gpu.size() / 2; 
      
      const int DIVISOR = 32767;

      std::cout << std::endl << "loadValues to TSDF in CUDA. We have  " << numPoints << " points to load" << std::endl;

      const short2 *startPtrConst = &(volume.ptr (0)[0]);
      short2  *startPtr = const_cast<short2*>(startPtrConst);
      
      //std::cout << "First data pointer: " << startPtr << std::endl;

      const int BlockX = 512;
      const int BlockY = 1;
      const int BlockSize = BlockX * BlockY;
      int numBlocks = divUp(numPoints, BlockSize);
      dim3 threadsPerBlock(BlockX, BlockY);
      
      cudaPrintfInit();
      
      loadTsdfValueKernel <<< numBlocks, threadsPerBlock >>> (startPtr, cloud_gpu.ptr(), numPoints, DIVISOR);
      cudaSafeCall ( cudaGetLastError () );
      cudaSafeCall (cudaDeviceSynchronize ());
        
      cudaPrintfDisplay(stdout, true);
      cudaPrintfEnd();
      
    }
    
  }  
}

